import triton
import triton.language as tl

import torch
import torch_npu
from .util import triton_helper_compile, compiled_kernel_list

from torch._inductor.runtime.triton_helpers import math as tl_math

# triton_poi_fused_add_cat_mul_2
@triton.jit
def fused_kernel9(in_ptr0, in_ptr1, out_ptr0, xnumel, XBLOCK : tl.constexpr):
    xnumel = 2048
    xoffset = tl.program_id(0) * XBLOCK
    xindex = xoffset + tl.arange(0, XBLOCK)[:]
    xmask = tl.full([XBLOCK], True, tl.int1)
    x2 = xindex
    x0 = xindex % 64
    x1 = (xindex // 64)
    tmp0 = tl.load(in_ptr0 + (x2), None)
    tmp1 = tl.load(in_ptr1 + (x2 % 32), None, eviction_policy='evict_last')
    tmp2 = tl_math.cos(tmp1)
    tmp3 = 1.0
    tmp4 = tmp2 * tmp3
    tmp5 = tmp0 * tmp4
    tmp6 = x0
    tmp7 = tl.full([1], 0, tl.int32)
    tmp8 = tmp6 >= tmp7
    tmp9 = tl.full([1], 32, tl.int32)
    tmp10 = tmp6 < tmp9
    tmp11 = tl.load(in_ptr0 + (32 + (64*x1) + x0), tmp10, eviction_policy='evict_last', other=0.0)
    tmp12 = -tmp11
    tmp13 = tl.full(tmp12.shape, 0.0, tmp12.dtype)
    tmp14 = tl.where(tmp10, tmp12, tmp13)
    tmp15 = tmp6 >= tmp9
    tmp16 = tl.full([1], 64, tl.int32)
    tmp17 = tmp6 < tmp16
    tmp18 = tl.load(in_ptr0 + ((64*x1) + ((-32) + x0)), tmp15, eviction_policy='evict_last', other=0.0)
    tmp19 = tl.where(tmp10, tmp14, tmp18)
    tmp20 = tl_math.sin(tmp1)
    tmp21 = tmp20 * tmp3
    tmp22 = tmp19 * tmp21
    tmp23 = tmp5 + tmp22
    tmp24 = 0.3535533905932738
    tmp25 = tmp23 * tmp24
    tl.store(out_ptr0 + (x2), tmp25, None)

def test_compile_for_fused_kernel9():
    signature = {0: '*fp32', 1: '*fp32', 2: '*fp32', 3: 'i32'}

    constants = {
        "XBLOCK": 2048
    }
    triton_helper_compile(func=fused_kernel9, signature=signature, constants=constants)

compiled_kernel_list["fused_kernel9"] = test_compile_for_fused_kernel9